/****************************************************************************************
*                                                                                      *
*                          Projeto BioInfo - Lincs and CIn / UFPE                      *
*                                    07/10/2014                                        *
*																					   *
****************************************************************************************
* Descricao: Arquivo que contem funcoes da classe Ocl.                    *
****************************************************************************************
* Responsaveis: Jefferson Ramos(jrla)													   *
*                                                                                      */
/*                                                                                     *
****************************************************************************************/
#include <Ocl.hpp>


int Ocl::shutdown(){

	// release resources
	for (unsigned int dev = 0; dev < num_devices; dev++)
	if (cmd_queue[dev]) clReleaseCommandQueue(cmd_queue[dev]); else return -1;
	if (context) clReleaseContext(context); else return -1;
	if (device_list) free(device_list); else return -1;


	clReleaseMemObject(score_vector_mem);
	clReleaseMemObject(query_mem);
	clReleaseMemObject(bank_mem);
	clReleaseMemObject(taxa_mem);

	// reset all variables
	context = 0;
	device_list = 0;
	num_devices = 0;
	device_type = 0;

	return 0;
}


void Ocl::usage(int argc, char **argv)
{
	fprintf(stderr, "usage: %s <max_rows/max_cols> <penalty> \n", argv[0]);
	fprintf(stderr, "\t<dimension>  - x and y dimensions\n");
	fprintf(stderr, "\t<penalty> - penalty(positive integer)\n");
	fprintf(stderr, "\t<file> - filename\n");
	exit(1);
}


int Ocl::initializeCL(int use_gpu)
{
	cl_uint num_platforms;
	cl_platform_id platform_id;
	char buffer[50];

	// create OpenCL context
	if (clGetPlatformIDs(0, NULL, &num_platforms) != CL_SUCCESS)
	{
		printf("ERROR: clGetPlatformIDs(1,*,0) failed\n"); return -1;
	}

	if (num_platforms > 0){
		cl_platform_id* platforms = (cl_platform_id*)malloc(num_platforms* sizeof(cl_platform_id));
		clGetPlatformIDs(num_platforms, platforms, NULL);

		platform_id = platforms[0];
		free(platforms);

	}
	else { printf("ERROR: num_platforms == 0 failed\n"); return -1; }

	//	printf("platform_id %lX num_platforms %d\n",platform_id,num_platforms);

	device_type = use_gpu ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU;//CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU;

	//	printf("use_gpu %d device_type %d\n",use_gpu, device_type);

	clGetDeviceIDs(platform_id, device_type, 0, NULL, &num_devices);

	device_list = (cl_device_id*)malloc(num_devices * sizeof(cl_device_id));
	if (device_list == NULL){ printf("ERROR: device_list == NULL\n"); return -1; }

	clGetDeviceIDs(platform_id, device_type, num_devices, device_list, NULL);

	//	printf("num_devices %d device_id %d\n", num_devices, device_list[0]);

	context = clCreateContext(NULL, num_devices, device_list, NULL, NULL, NULL);

	cmd_queue = (cl_command_queue*)malloc(num_devices * sizeof(cl_command_queue));

	for (unsigned int dev = 0; dev < num_devices; dev++)
		cmd_queue[dev] = clCreateCommandQueue(context, device_list[dev], 0, NULL);

	clGetDeviceInfo(device_list[0], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(buf_uint), &buf_uint, NULL);

	//	printf("\nMAX: %lf Megas\n",(double)buf_uint/(1024*1024));

	clGetDeviceInfo(device_list[0], CL_DEVICE_NAME, sizeof(buffer), buffer, NULL);

	//	printf("DEVICE_NAME = %s\nDEVICE_MAX_MEM_ALLOC_SIZE = %f Mb\n", buffer,(float)buf_uint/(1024*1024));

	return 0;
}


int Ocl::initializeKernel()
{
	FILE * log;
	char buildLog[16384];
	int sourcesize = 1024 * 1024;

	char * source = (char *)calloc(sourcesize, sizeof(char));
	if (!source) { printf("ERROR: calloc(%d) failed\n", sourcesize); return -1; }

	// read the kernel core source
	FILE * fp;
	fopen_s(&fp, "../OCL/src/nw.cl", "rb");
	if (!fp) { printf("ERROR: unable to open\n"); return -1; }

	fread(source + strlen(source), sourcesize, 1, fp);
	fclose(fp);

	// compile kernel
	const char * slist[2] = { source, 0 };
	size_t sourceSize[] = { strlen(source) };

	cl_program prog = clCreateProgramWithSource(context,
		1,
		slist,
		sourceSize,
		&err);
	if (err != CL_SUCCESS) { printf("ERROR: clCreateProgramWithSource() => %d\n", err); return -1; }

	err = clBuildProgram(prog, 1, device_list, NULL, NULL, NULL);

	// Determine the reason for the error
	clGetProgramBuildInfo(prog,
		device_list[0],
		CL_PROGRAM_BUILD_LOG,
		sizeof(buildLog),
		buildLog,
		NULL);

	if (err != CL_SUCCESS){
		printf("ERROR: clBuildProgram() => %d\n buildLog: \n %s", err, buildLog); return -1;
	}
	else{
		fopen_s(&log, "buildLog.txt", "w");
		fprintf(log, "\nINFO: clBuildProgram() => %d\n buildLog: \n %s", err, buildLog);
	}

	kernel1 = clCreateKernel(prog, "nw_kernel1", &err);

	if (err != CL_SUCCESS) { printf("ERROR: clCreateKernel() 0 => %d\n", err); return -1; }
	clReleaseProgram(prog);

	return 0;
}


int Ocl::initializeCLBuffers()
{
	score_vector_mem = clCreateBuffer(context, CL_MEM_READ_WRITE, (TAMSEQ + 1)* sizeof(int)*numIndividuals_locationMaximum*alignment_by_individuals, NULL, &err);
	if (err != CL_SUCCESS) { printf("ERROR: clCreateBuffer reference score_vector_mem(size:%d) => %d\n", (TAMSEQ + 1)*numIndividuals_locationMaximum*alignment_by_individuals, err); return -1; }

	query_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(char)*TAMSEQ*alignment_by_individuals, NULL, &err);
	if (err != CL_SUCCESS) { printf("ERROR: clCreateBuffer reference query_mem(size:%d) => %d\n", TAMSEQ*alignment_by_individuals, err); return -1; }

	bank_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, TAMSEQ*sizeof(char)*alignment_by_individuals*numIndividuals_locationMaximum, NULL, &err);
	if (err != CL_SUCCESS) { printf("ERROR: clCreateBuffer reference bank_mem(size:%d) => %d\n", TAMSEQ *alignment_by_individuals*num_individuals, err); return -1; }

	taxa_mem = clCreateBuffer(context, CL_MEM_READ_WRITE, num_individuals*alignment_by_individuals*sizeof(int), NULL, &err);
	if (err != CL_SUCCESS) { printf("ERROR: clCreateBuffer reference taxa_mem(size:%d) => %d\n", num_individuals*alignment_by_individuals, err); return -1; }

	err = clEnqueueWriteBuffer(cmd_queue[0], query_mem, 1, 0, TAMSEQ * sizeof(char)*alignment_by_individuals, query, 0, 0, 0);
	if (err != CL_SUCCESS) { printf("ERROR: clEnqueueWriteBuffer query (size:%d) => %d\n", TAMSEQ*alignment_by_individuals, err); return -1; }

	return 0;
}


void Ocl::setKernelArgs(){
	int Arg = 0;

	clSetKernelArg(kernel1, Arg++, sizeof(cl_mem), &query_mem);
	clSetKernelArg(kernel1, Arg++, sizeof(cl_mem), &score_vector_mem);
	clSetKernelArg(kernel1, Arg++, sizeof(cl_mem), &bank_mem);
	clSetKernelArg(kernel1, Arg++, sizeof(cl_mem), &taxa_mem);
	clSetKernelArg(kernel1, Arg++, sizeof(int), &alignment_by_individuals);
	clSetKernelArg(kernel1, Arg++, sizeof(int), &num_individuals);
}


int Ocl::runCLKernels()
{
	int i;
	int numWorkGroups;
	size_t localWorkSize = 1;
	size_t globalWorkSize;

	if (split){
		for (i = 0; i < individuals_index - 1; i++){
			err = clEnqueueWriteBuffer(cmd_queue[0], bank_mem, CL_TRUE, 0, TAMSEQ*alignment_by_individuals*numIndividuals_locationMaximum*sizeof(char), individuals[i], 0, 0, 0);
			if (err != CL_SUCCESS) { printf("ERROR: clEnqueueWriteBuffer individuals (size:%d) => %d\n", TAMSEQ*alignment_by_individuals*numIndividuals_locationMaximum, err); return -1; }

			numWorkGroups = ((alignment_by_individuals*numIndividuals_locationMaximum) + localWorkSize - 1) / localWorkSize;// round up
			globalWorkSize = numWorkGroups * localWorkSize;// must be evenly divisible by

			err = clEnqueueNDRangeKernel(cmd_queue[0], kernel1, 1, NULL, &globalWorkSize, &localWorkSize, 0, NULL, NULL);
			if (err != CL_SUCCESS) { printf("ERROR: 1  clEnqueueNDRangeKernel()=>%d failed\n", err); return -1; }

			err = clEnqueueReadBuffer(cmd_queue[0], taxa_mem, CL_TRUE, 0, numIndividuals_locationMaximum*alignment_by_individuals*sizeof(int), &output_taxa[numIndividuals_locationMaximum*alignment_by_individuals*i], 0, 0, 0);
		}

		if ((num_individuals - numIndividuals_locationMaximum*i) > 0){
			err = clEnqueueWriteBuffer(cmd_queue[0], bank_mem, CL_TRUE, 0, TAMSEQ * sizeof(char)*alignment_by_individuals*(num_individuals - numIndividuals_locationMaximum*i), individuals[i], 0, 0, 0);
			if (err != CL_SUCCESS) { printf("ERROR: clEnqueueWriteBuffer individuals (size:%d) => %d\n", TAMSEQ*alignment_by_individuals*(num_individuals - numIndividuals_locationMaximum*i), err); return -1; }

			numWorkGroups = ((alignment_by_individuals*(num_individuals - numIndividuals_locationMaximum*i)) + localWorkSize - 1) / localWorkSize;// round up
			globalWorkSize = numWorkGroups * localWorkSize;// must be evenly divisible by

			err = clEnqueueNDRangeKernel(cmd_queue[0], kernel1, 1, NULL, &globalWorkSize, &localWorkSize, 0, NULL, NULL);
			if (err != CL_SUCCESS) { printf("ERROR: 1  clEnqueueNDRangeKernel()=>%d failed\n", err); return -1; }

			err = clEnqueueReadBuffer(cmd_queue[0], taxa_mem, CL_TRUE, 0, (num_individuals - numIndividuals_locationMaximum*i)*alignment_by_individuals*sizeof(int), &output_taxa[numIndividuals_locationMaximum*alignment_by_individuals*i], 0, 0, 0);
		}
	}
	else{
		err = clEnqueueWriteBuffer(cmd_queue[0], bank_mem, CL_TRUE, 0, TAMSEQ * sizeof(char)*alignment_by_individuals*num_individuals, individuals[0], 0, 0, 0);
		if (err != CL_SUCCESS) { printf("ERROR: clEnqueueWriteBuffer individuals (size:%d) => %d\n", TAMSEQ*alignment_by_individuals*num_individuals, err); return -1; }

		numWorkGroups = ((alignment_by_individuals*num_individuals) + localWorkSize - 1) / localWorkSize;// round up
		globalWorkSize = numWorkGroups * localWorkSize;// must be evenly divisible by

		err = clEnqueueNDRangeKernel(cmd_queue[0], kernel1, 1, NULL, &globalWorkSize, &localWorkSize, 0, NULL, NULL);
		if (err != CL_SUCCESS) { printf("ERROR: 1  clEnqueueNDRangeKernel()=>%d failed\n", err); return -1; }

		err = clEnqueueReadBuffer(cmd_queue[0], taxa_mem, CL_TRUE, 0, num_individuals*alignment_by_individuals*sizeof(int), output_taxa, 0, 0, 0);
	}
	return 0;
}


int Ocl::vectors_malloc()
{
	int i;

	numIndividuals_locationMaximum = 500;
	individuals_index = (int)ceil((float)num_individuals / numIndividuals_locationMaximum);//quantidade de divisoes
	split = num_individuals > numIndividuals_locationMaximum ? 1 : 0;

	//	printf("numIndividuals_locationMaximum:%d  -   %d\n",numIndividuals_locationMaximum,split);

	query = (char*)_aligned_malloc(sizeof(char)*TAMSEQ*alignment_by_individuals, 64);
	if (!query){ printf("ERROR: Alloc_query(size:%d)\n", TAMSEQ*alignment_by_individuals); return -1; }
	output_taxa = (int *)_aligned_malloc(alignment_by_individuals*num_individuals*sizeof(int), 64);
	if (!output_taxa){ printf("ERROR: Alloc_output_taxa(size:%d)\n", alignment_by_individuals*num_individuals); return -1; }
	store_individuals = (char *)_aligned_malloc((sizeof(char)*(TAMSEQ / Byte_By_TwoBit)*alignment_by_individuals*num_individuals), 64);
	if (!store_individuals){ printf("ERROR: Alloc_store_individuals(size:%u)\n", ((TAMSEQ / Byte_By_TwoBit)*alignment_by_individuals*num_individuals)); return -1; }
	individuals = (char**)_aligned_malloc(sizeof(char*)*individuals_index, 64);
	if (!individuals){ printf("ERROR: Alloc_Individuals**(size:%d)\n", individuals_index); return -1; }

	if (split){
		for (i = 0; i < individuals_index - 1; i++){
			individuals[i] = (char*)_aligned_malloc(sizeof(char)*TAMSEQ*alignment_by_individuals*numIndividuals_locationMaximum, 64);
			if (!individuals[i]){ printf("ERROR: Alloc_Individuals*(size:%d)\n", TAMSEQ*alignment_by_individuals*numIndividuals_locationMaximum); return -1; }
		}

		if ((num_individuals - numIndividuals_locationMaximum*i) > 0){
			individuals[i] = (char*)_aligned_malloc(sizeof(char)*TAMSEQ*alignment_by_individuals*(num_individuals - numIndividuals_locationMaximum), 64);
			if (!individuals[i]){ printf("ERROR: Alloc_Individuals*(size:%d)\n", TAMSEQ*alignment_by_individuals*(num_individuals - numIndividuals_locationMaximum*i)); return -1; }
		}
		//printf("\nNum_Division: %d\n", individuals_index);
	}
	else{
		individuals[0] = (char*)_aligned_malloc(sizeof(char)*TAMSEQ*alignment_by_individuals*num_individuals, 64);
		if (!individuals[0]){ printf("ERROR: Alloc_Individuals*(size:%d)\n", TAMSEQ*alignment_by_individuals*num_individuals); return -1; }
		numIndividuals_locationMaximum = num_individuals;
	}
	return 0;
}


void Ocl::read_input_individuals() {

	int i, j, k;

	if (split){
		for (i = 0; i < individuals_index - 1; i++)
		{
			fread(store_individuals, sizeof(char), (TAMSEQ / Byte_By_TwoBit)*alignment_by_individuals*numIndividuals_locationMaximum, BD_nucleotideos);

			for (j = 0, k = 0; j < (TAMSEQ / Byte_By_TwoBit)*alignment_by_individuals*numIndividuals_locationMaximum; j++){

				individuals[i][k++] = (MASK_NCL1 & store_individuals[j]) + 48;
				individuals[i][k++] = ((MASK_NCL2 & store_individuals[j]) >> 2) + 48;
				individuals[i][k++] = ((MASK_NCL3 & store_individuals[j]) >> 4) + 48;
				individuals[i][k++] = ((MASK_NCL4 & store_individuals[j]) >> 6) + 48;
			}
		}

		if (num_individuals - numIndividuals_locationMaximum*i > 0){

			fread(store_individuals, sizeof(char), (TAMSEQ / Byte_By_TwoBit)*alignment_by_individuals*(num_individuals - numIndividuals_locationMaximum*i), BD_nucleotideos);

			for (j = 0, k = 0; j < (TAMSEQ / Byte_By_TwoBit)*alignment_by_individuals*(num_individuals - numIndividuals_locationMaximum*i); j++){

				individuals[i][k++] = (MASK_NCL1 & store_individuals[j]) + 48;
				individuals[i][k++] = ((MASK_NCL2 & store_individuals[j]) >> 2) + 48;
				individuals[i][k++] = ((MASK_NCL3 & store_individuals[j]) >> 4) + 48;
				individuals[i][k++] = ((MASK_NCL4 & store_individuals[j]) >> 6) + 48;
			}
		}

	}
	else{

		fread(store_individuals, sizeof(char), (TAMSEQ / Byte_By_TwoBit)*alignment_by_individuals*num_individuals, BD_nucleotideos);

		for (j = 0, k = 0; j < (TAMSEQ / Byte_By_TwoBit)*alignment_by_individuals*num_individuals; j++){

			individuals[0][k++] = (MASK_NCL1 & store_individuals[j]) + 48;
			individuals[0][k++] = ((MASK_NCL2 & store_individuals[j]) >> 2) + 48;
			individuals[0][k++] = ((MASK_NCL3 & store_individuals[j]) >> 4) + 48;
			individuals[0][k++] = ((MASK_NCL4 & store_individuals[j]) >> 6) + 48;
		}
	}
}


void Ocl::read_input_query() {

	int j, k = 0;
	char nucleotideo;

	fj(0, (TAMSEQ / Byte_By_TwoBit)*alignment_by_individuals*query_compare){
		fread(&nucleotideo, sizeof(char), 1, BD_nucleotideos);
		if (j >= ((TAMSEQ / Byte_By_TwoBit)*alignment_by_individuals*(query_compare - 1))){
			query[k++] = (MASK_NCL1 & nucleotideo) + 48;
			query[k++] = ((MASK_NCL2 & nucleotideo) >> 2) + 48;
			query[k++] = ((MASK_NCL3 & nucleotideo) >> 4) + 48;
			query[k++] = ((MASK_NCL4 & nucleotideo) >> 6) + 48;
		}
	}
	query[k] = '\0';
}


void Ocl::print_results(){

	for (int i = 1, taxai = 0; i <= num_individuals*alignment_by_individuals; i++){
		taxai += output_taxa[i - 1];

		if (i % (alignment_by_individuals) == 0){
			fprintf(TaxaOut, "%d\n", taxai);
			taxai = 0;
		}
	}
}
